#include <stdio.h>
#include <stdlib.h>
#include <cstring>
// includes, project
#include "cutil_inline.h"
#include "DPComparator.h"
#include "Matrix.h"
#include <iostream> 
//#include "SURFeaturesCUDA.cuh"

//#include "cuPrintf.cuh"
//extern "C" void euclideanDistance(float* descriptorsObject,size_t descriptorsObjectSize,float* descriptorImage, size_t descriptorImageSize);

__global__ void euclideanDistanceKernal(float* descriptorsObject,int noOfKeyPointsObject,float* descriptorImage, int noOfKeyPointsImage,float* ObjectImageDistance )  
{
  //  int idx = blockIdx.x * blockDim.x + threadIdx.x+1;
   /* float* discriptorObjectItem;
    float* discriptorImageItem;*/

   /* int row = blockIdx.y * blockDim.y + threadIdx.y;      
	int col = blockIdx.x * blockDim.x + threadIdx.x;*/
	
	int row = blockIdx.x * blockDim.x + threadIdx.x; 
	int col = blockIdx.y * blockDim.y + threadIdx.y;
	
	/*int col = blockIdx.y * blockDim.y + threadIdx.y;      
	int row = blockIdx.x * blockDim.x + threadIdx.x;*/
   // int row = threadIdx.y; 
	//int col = threadIdx.x; 
	//printf("From the GPU idx :%d\n",idx);
	 float eiuclidianDistance=-1;
	
	// printf("%d, %d: ",noOfKeyPointsObject,noOfKeyPointsImage);
	 if(row<noOfKeyPointsObject&&col<noOfKeyPointsImage)
	 {
	   // if(col>239&&row>240)
	   // {
	       // printf("%d, %d: ",row,col);
        eiuclidianDistance=0;
        for(int i=0;i<128;i++)
        {
            float diff = descriptorsObject[row*128+i] - descriptorImage[col*128+i]; 
            eiuclidianDistance = eiuclidianDistance + diff*diff; 
            //eiuclidianDistance=eiuclidianDistance+ descriptorImage[col*128+i]*descriptorImage[col*128+i];
           //eiuclidianDistance= descriptorsObject[row*128+i]*descriptorsObject[row*128+i];
           // ObjectImageDistance[row*noOfKeyPointsImage+col] = diff*diff; 
            //ObjectImageDistance[row*noOfKeyPointsImage+col] = descriptorImage[col*128+i]*descriptorImage[col*128+i]; 
            // printf("From the GPU row: %d %d %.4f, col: %d %d %.4f\n",row,(row*128+i),descriptorsObject[row*128+i],col,(col*128+i),descriptorImage[col*128+i]);   	    
        }
	  //  }
	    ObjectImageDistance[row*noOfKeyPointsImage+col]= eiuclidianDistance;
	    
	   // cudaThreadSynchronize();
	   //ObjectImageDistance[row*noOfKeyPointsImage+col]= row*1000+col;
	    //printf("%.2f ",eiuclidianDistance);
	  // printf("From the GPU row:%d, col:%d\n",row,col);
	 }
	/*if(row*128<=descriptorsObjectSizeInt&&col*128<=descriptorImageSizeInt)
	//if(*128<=descriptorsObjectSizeInt&&col*128<=descriptorImageSizeInt)
	{   	   
	    eigenDistance=0;
	    for(int i=0;i<128;i++)
	    {   
           // if(row*128>30900&&col*128>30900)
           //if(row*128>30900&&col*128>30800)
           {
                printf("From the GPU row:%d %d %d %.4f, col:%d %d %d %.4f\n",descriptorsObjectSizeInt,row,(row*128+i),descriptorsObject[row*128+i],descriptorImageSizeInt,col,(col*128+i),descriptorImage[col*128+i]);
           }
	   

	    }
    	
	    {
	    //printf("From the GPU row:%d, col:%d\n",row,col);
	    }
	}*/
	
} 

__global__ void dpErrorKernal(int noOfKeyPointsObject, int noOfKeyPointsImage,float* ObjectImageDistance)
{
    int row = blockIdx.x * blockDim.x + threadIdx.x; 
	int col = blockIdx.y * blockDim.y + threadIdx.y;
	
	 if(row<noOfKeyPointsObject&&col<noOfKeyPointsImage)
	 {
	 
	 }


}   

__global__ void minErrorKernal(int noOfKeyPointsObject, int noOfKeyPointsImage,float* d_ObjectImageDistance,float* d_MinError,int* d_MinErrorLocation)
{
  /*    	int row = blockIdx.x * blockDim.x + threadIdx.x; 
	    int col = blockIdx.y * blockDim.y + threadIdx.y;*/
	    
	    int i = threadIdx.x; 
	    d_MinError[i] = 9999999999;
	  
         if(i<noOfKeyPointsObject)
         {
               for(int j=0;j<noOfKeyPointsImage;j++)
               {
    	           if(d_ObjectImageDistance[i*noOfKeyPointsObject+j]<d_MinError[j])
    	           {
    	             d_MinError[i]=d_ObjectImageDistance[i*noOfKeyPointsObject+j];
    	             d_MinErrorLocation[i]= j;
    	           }
               }
         }
         
        // printf("Min Error:%.2f, location:%d\n", d_MinError[i],d_MinErrorLocation[i]);
	       
}                                                                                                                               // 30848                                             //30976
//extern "C" void euclideanDistance(float* descriptorsObject,int descriptorsObjectSizeInt,float* descriptorImage, int descriptorImageSizeInt)
//                              (descriptorsMinArray,minSize,descriptorsMaxArray,maxSize);
extern "C" float getDPDistance(float* descriptorsMinArray,int minSize,float* descriptorsMaxArray, int maxSize)
//void euclideanDistance(float* descriptorsObject,size_t descriptorsObjectSize,float* descriptorImage, size_t descriptorImageSize)
//extern "C" void euclideanDistance()  extern "C" 
{
  //  int blockSize=32;
        
    int deviceCount = 0;

    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0)
    {
            printf("There are no available device(s) that support CUDA\n");
    }
    //cutilSafeCall(cudaDeviceReset());
    /*size_t free;  
    size_t total;  
    cuMemGetInfo(&free, &total);
    printf("free memory: %.2f, total memory: %.2f/n",(free / 1024 / 1024),(total / 1024 / 1024)); */
    
    size_t freeSize;  
    size_t totalSize;
    cuMemGetInfo_v2(&freeSize, &totalSize);
    printf("free memory: %.2f, total memory: %.2f/n",(float)freeSize/1024/1024,((float)totalSize / 1024 / 1024)); 
    //cout << "free memory: " << freeSize / 1024 / 1024 << "mb, total memory: " << totalSize / 1024 / 1024 << "mb" << endl;  
    
    float* d_descriptorsObject;
    float* d_descriptorImage;
    float* d_ObjectImageDistance;
    
    
    
    
    size_t descriptorsObjectSize=sizeof(float)*minSize;
    size_t descriptorImageSize=sizeof(float)*maxSize;
    
    
    cutilSafeCall(cudaMalloc(&d_descriptorsObject, descriptorsObjectSize));
    cutilSafeCall(cudaMalloc(&d_descriptorImage, descriptorImageSize));
    
    cutilSafeCall(cudaMemcpy(d_descriptorsObject, descriptorsMinArray, descriptorsObjectSize, cudaMemcpyHostToDevice));
    cutilSafeCall(cudaMemcpy(d_descriptorImage, descriptorsMaxArray, descriptorImageSize, cudaMemcpyHostToDevice));
    
    int noOfKeyPointsObject = minSize/128;
    int noOfKeyPointsImage = maxSize/128;
    
    printf("Points Object:%d, Points Image:%d\n",noOfKeyPointsObject,noOfKeyPointsImage);
    
    size_t  objectImageDistanceSize = sizeof(float)*noOfKeyPointsObject*noOfKeyPointsImage;
    cutilSafeCall(cudaMalloc(&d_ObjectImageDistance, objectImageDistanceSize));
	
	
	
    
	//dim3 dimBlock(blockSize,blockSize); 
	
   // int N = max(descriptorsObjectSizeInt,descriptorImageSizeInt);
    int block_size = 32;

  // int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); 
//   primeCounter<<< n_blocks, block_size >>> (primeCount_d, N);

    dim3 dimBlock(block_size,block_size); 
    float devider = block_size;
    
    int gridSizeObject = 0;
    int gridSizeImage = 0;
    if(noOfKeyPointsObject%block_size==0){gridSizeObject = noOfKeyPointsObject/block_size + 1;}else{gridSizeObject=ceil(noOfKeyPointsObject/(float)block_size);}
    if(noOfKeyPointsImage%block_size==0){gridSizeImage = noOfKeyPointsImage/block_size + 1;}else{gridSizeImage=ceil(noOfKeyPointsImage/(float)block_size);}
    
    printf("%d,%d : %d,%d\n",noOfKeyPointsObject,noOfKeyPointsImage,gridSizeObject,gridSizeImage);     
    //printf("%.4f,%.4f\n",descriptorsObjectSizeInt/devider,descriptorImageSizeInt/devider);
    dim3 dimGrid(ceil(noOfKeyPointsObject/devider),ceil(noOfKeyPointsImage/devider));
   // dim3 dimGrid(gridSizeObject,gridSizeImage);
//dim3 dimGrid(ceil(width/(float)blockSize),ceil(height/(float)blockSize),noOfImages);
    
    //euclideanDistanceKernal<<<dimGrid, dimBlock>>>(d_descriptorsObject,d_descriptorImage);
    //euclideanDistanceKernal<<<n_blocks, block_size>>>(d_descriptorsObject,descriptorsObjectSizeInt,d_descriptorImage,descriptorImageSizeInt);
    euclideanDistanceKernal<<<dimGrid, dimBlock>>>(d_descriptorsObject,noOfKeyPointsObject,d_descriptorImage,noOfKeyPointsImage,d_ObjectImageDistance);
    cutilSafeCall(cudaThreadSynchronize());
    
    //float* h_ObjectImageDistance = (float*) malloc(objectImageDistanceSize); //sizeof(float)
    /*float* h_ObjectImageDistance = (float*) malloc(objectImageDistanceSize); //sizeof(float)
    cutilSafeCall(cudaMemcpy(h_ObjectImageDistance, d_ObjectImageDistance, objectImageDistanceSize, cudaMemcpyDeviceToHost)); */
      
      
     cutilSafeCall(cudaFree(d_descriptorsObject));
     cutilSafeCall(cudaFree(d_descriptorImage)); 
              
     float* d_MinError;
     int* d_MinErrorLocation;
     
     size_t  minErrorSize = sizeof(float)*noOfKeyPointsObject;
     size_t  minErrorLocationSize = sizeof(int)*noOfKeyPointsObject;
     cutilSafeCall(cudaMalloc(&d_MinError, minErrorSize));
     cutilSafeCall(cudaMalloc(&d_MinErrorLocation, minErrorLocationSize));
     
     minErrorKernal<<<1, noOfKeyPointsObject>>>(noOfKeyPointsObject,noOfKeyPointsImage,d_ObjectImageDistance, d_MinError,d_MinErrorLocation);
     cutilSafeCall(cudaThreadSynchronize()); 
     
    /* float* d_descriptorsObject;
     float* d_errorOutput; 
     
     dpErrorKernal<<<dimGrid, dimBlock>>>(noOfKeyPointsObject,d_ObjectImageDistance);     */  
     cutilSafeCall(cudaFree(d_ObjectImageDistance));
     
     float* h_MinError = (float*) malloc(minErrorSize); //sizeof(float)
    cutilSafeCall(cudaMemcpy(h_MinError, d_MinError, minErrorSize, cudaMemcpyDeviceToHost));
    
     int* h_MinErrorLocation = (int*) malloc(minErrorLocationSize); //sizeof(float)
    cutilSafeCall(cudaMemcpy(h_MinErrorLocation, d_MinErrorLocation, minErrorLocationSize, cudaMemcpyDeviceToHost)); 
     
     cutilSafeCall(cudaFree(d_MinError));
     cutilSafeCall(cudaFree(d_MinErrorLocation));
     
     float totalError=0;
     for(int j=0;j<noOfKeyPointsObject;j++)   
     {
        totalError = totalError + h_MinError[j];            
     }
     printf("Total average Error:%.2f ",totalError/noOfKeyPointsObject);
    // cuMemGetInfo_v2(&freeSize, &totalSize);
     //printf("free memory: %.2f, total memory: %.2f/n",(float)freeSize/1024/1024,((float)totalSize / 1024 / 1024)); 
   
     /*for(int i=0;i<noOfKeyPointsObject;i++)
     {
        for(int j=0;j<noOfKeyPointsImage;j++)   
        {
            printf("%.2f ",h_ObjectImageDistance[i*noOfKeyPointsImage+j]);
        }
        printf("\n");
     }*/
    /*int row = blockIdx.y * blockDim.y + threadIdx.y; 
	int col = blockIdx.x * blockDim.x + threadIdx.x;
	
	const DevMem2Db& query, const DevMem2Db& train
    cv::gpu::GpuMat diffDescriptors;
     diffDescriptors= objectDescriptorsGpuMat-imageDescriptorsGpuMat;*/
     
    /* float** tmpMat = new float*[noOfKeyPointsObject];
	 float* tmpArr;
	
     for(int i=0;i<noOfKeyPointsObject;i++)
     {
        tmpArr = new float[noOfKeyPointsImage];
        memcpy(tmpArr,&h_ObjectImageDistance[i*noOfKeyPointsImage],sizeof(float)*noOfKeyPointsImage);
        tmpMat[i]=tmpArr;
     }
     free(h_ObjectImageDistance);
     
     Matrix* errorMatrix  = new Matrix(tmpMat,noOfKeyPointsObject,noOfKeyPointsImage);
     
     DPComparator* dpComparator = new  DPComparator();
     dpComparator->setErrorMatrix(errorMatrix,noOfKeyPointsObject,noOfKeyPointsImage);
    // dpComparator->setErrorMatrix(errorMatrix,noOfKeyPointsImage,noOfKeyPointsObject);
    // dpComparator->setMinMaxSize();
   // float returnFloat=10;
     float returnFloat = dpComparator->calculateDPError();
     delete  dpComparator;
     */
       
      // free(h_ObjectImageDistance);
     float returnFloat=10;

     
     return returnFloat; 
}

